/**
   device_controller.cu - Created by Timothy Morey on 10/10/2011.
*/

#include <time.h>

#include "device_controller.h"
#include "parts.cuh"
#include "simulator.h"

DeviceController::DeviceController(unsigned char did, 
				   unsigned int minPart,
				   unsigned int maxPart,
                                   pthread_barrier_t* barrier)
  : _DeviceID(did),
    _MinParticleID(minPart),
    _MaxParticleID(maxPart),
    _CurStep(0),
    _DevX(0),
    _DevY(0),
    _DevVX(0),
    _DevVY(0),
    _DevVM(0),
    _DevAX(0),
    _DevAY(0),
    _AverageGPUTime(0.0),
    _AverageCPUInnerTime(0.0),
    _AverageCPUOuterTime(0.0),
    _GPUElapsedTime(0.0),
    _Barrier(barrier)
{
  pthread_create(&_ThreadID, 0, DeviceController::ThreadEntryPoint, this);
}

DeviceController::~DeviceController()
{
  pthread_join(_ThreadID, 0);
}

void* DeviceController::ThreadEntryPoint(void* arg)
{
  ((DeviceController*)arg)->ThreadLoop();
  return 0;
}

void DeviceController::ThreadLoop()
{
  CUDA_CALL( cudaSetDevice(_DeviceID) );
  this->InitDevMemory();

  double elapsedOuterSum = 0.0;
  double elapsedInnerSum = 0.0;

  timespec innerStart, innerEnd, outerStart, outerEnd;
  clock_gettime(CLOCK_MONOTONIC, &outerStart);

  while(true)
  {
    if(Simulator::Instance().StopFlag())
    {
      break;
    }
    else if(Simulator::Instance().PauseFlag())
    {
      pthread_yield();
    }
    else
    {
      pthread_barrier_wait(_Barrier);

      clock_gettime(CLOCK_MONOTONIC, &innerStart);
      this->Step();
      clock_gettime(CLOCK_MONOTONIC, &innerEnd);
      clock_gettime(CLOCK_MONOTONIC, &outerEnd);
      elapsedInnerSum += ((double)(innerEnd.tv_sec - innerStart.tv_sec)) +
        ((double)(innerEnd.tv_nsec - innerStart.tv_nsec)) / 1000000000.0;
      elapsedOuterSum += ((double)(outerEnd.tv_sec - outerStart.tv_sec)) +
        ((double)(outerEnd.tv_nsec - outerStart.tv_nsec)) / 1000000000.0;
      outerStart = outerEnd;
      _CurStep++;

      if(_CurStep % 100 == 0)
      {
        _AverageCPUInnerTime = elapsedInnerSum / 100.0;
        _AverageCPUOuterTime = elapsedOuterSum / 100.0;
        elapsedInnerSum = 0.0;
        elapsedOuterSum = 0.0;
      }
    }
  }

  this->FreeDevMemory();
}

void DeviceController::Step()
{
  cudaEvent_t start, stop;
  float elapsed;

  CUDA_CALL( cudaEventCreate(&start) );
  CUDA_CALL( cudaEventCreate(&stop) );
  CUDA_CALL( cudaEventRecord(start, 0) );

  Simulator& sim = Simulator::Instance();
  const Parameters& parms = sim.GetParameters();
  int partCount = _MaxParticleID - _MinParticleID + 1;

  if(_MinParticleID > 0)
  {
    // copy parts with id < _MinParticleID to device
    size_t memsize = _MinParticleID * sizeof(DoF);
    CUDA_CALL( cudaMemcpy(_DevX, sim.GetX(), memsize, cudaMemcpyHostToDevice) );
    CUDA_CALL( cudaMemcpy(_DevY, sim.GetY(), memsize, cudaMemcpyHostToDevice) );
  }

  if(_MaxParticleID < parms.N - 1)
  {
    // copy parts with id > _MaxParticleID to device
    size_t memsize = (parms.N - _MaxParticleID - 1) * sizeof(DoF);
    int offset = _MaxParticleID + 1;
    CUDA_CALL( cudaMemcpy(_DevX + offset, sim.GetX() + offset, 
                          memsize, cudaMemcpyHostToDevice) );
    CUDA_CALL( cudaMemcpy(_DevY + offset, sim.GetY() + offset, 
                          memsize, cudaMemcpyHostToDevice) );
  }
  
  UpdateAcceleration
    <<<partCount, min(THREADS_PER_BLOCK, parms.N)>>>
    (_DevX, _DevY, _DevAX, _DevAY, 
     parms.N, parms.MinX, parms.MaxX, parms.MinY, parms.MaxY,
     _MinParticleID, _MaxParticleID);

  UpdateVelocity
    <<<partCount, 1>>>
    (_DevVX, _DevVY, _DevAX, _DevAY, 
     parms.dt, _MinParticleID, _MaxParticleID);

  UpdatePositionAndVelocity 
    <<<partCount, 1>>> 
    (_DevX, _DevY, _DevVX, _DevVY, _DevAX, _DevAY, 
     parms.N, parms.dt, parms.MinX, parms.MaxX, parms.MinY, parms.MaxY,
     _MinParticleID, _MaxParticleID);

  CalculateMagnitude
    <<<partCount, 1>>>
    (_DevVX, _DevVY, _DevVM, _MinParticleID, _MaxParticleID);

  // copy local parts to host
  size_t memsize = partCount * sizeof(DoF);
  CUDA_CALL( cudaMemcpy(sim.GetX() + _MinParticleID, _DevX + _MinParticleID,
                        memsize, cudaMemcpyDeviceToHost) );
  CUDA_CALL( cudaMemcpy(sim.GetY() + _MinParticleID, _DevY + _MinParticleID,
                        memsize, cudaMemcpyDeviceToHost) );
  CUDA_CALL( cudaMemcpy(sim.GetVM() + _MinParticleID, _DevVM,
                        memsize, cudaMemcpyDeviceToHost) );

  CUDA_CALL( cudaEventRecord(stop, 0) );
  CUDA_CALL( cudaEventSynchronize(stop) );
  CUDA_CALL( cudaEventElapsedTime(&elapsed, start, stop) );
  CUDA_CALL( cudaEventDestroy(start) );
  CUDA_CALL( cudaEventDestroy(stop) );

  _GPUElapsedTime += elapsed;

  if(_CurStep % 100 == 0)
  {
    _AverageGPUTime = _GPUElapsedTime / 100.0;
    _GPUElapsedTime = 0.0;
  }
}

void DeviceController::InitDevMemory()
{
  size_t posSize = Simulator::Instance().GetParameters().N * sizeof(DoF);
  size_t localSize = (_MaxParticleID - _MinParticleID + 1) * sizeof(DoF);

  CUDA_CALL( cudaMalloc(&_DevX, posSize) );
  CUDA_CALL( cudaMalloc(&_DevY, posSize) );
  CUDA_CALL( cudaMalloc(&_DevVX, localSize) );
  CUDA_CALL( cudaMalloc(&_DevVY, localSize) );
  CUDA_CALL( cudaMalloc(&_DevVM, localSize) );
  CUDA_CALL( cudaMalloc(&_DevAX, localSize) );
  CUDA_CALL( cudaMalloc(&_DevAY, localSize) );

  this->SyncPartsHostToDevice();
}

void DeviceController::FreeDevMemory()
{
  this->SyncPartsDeviceToHost();

  if(_DevX) 
  {
    CUDA_CALL( cudaFree(_DevX) );
    _DevX = 0;
  }

  if(_DevY) 
  {
    CUDA_CALL( cudaFree(_DevY) );
    _DevY = 0;
  }

  if(_DevVX) 
  {
    CUDA_CALL( cudaFree(_DevVX) );
    _DevVX = 0;
  }
  
  if(_DevVY) 
  {
    CUDA_CALL( cudaFree(_DevVY) );
    _DevVY = 0;
  }
  
  if(_DevVM) 
  {
    CUDA_CALL( cudaFree(_DevVM) );
    _DevVM = 0;
  }

  if(_DevAX) 
  {
    CUDA_CALL( cudaFree(_DevAX) );
    _DevAX = 0;
  }

  if(_DevAY) 
  {
    CUDA_CALL( cudaFree(_DevAY) );
    _DevAY = 0;
  }
}

void DeviceController::SyncPartsHostToDevice()
{
  this->SyncPosHostToDevice();

  Simulator& sim = Simulator::Instance();
  size_t memsize = (_MaxParticleID - _MinParticleID + 1) * sizeof(DoF);

  CUDA_CALL( cudaMemcpy(_DevVX, sim.GetVX() + _MinParticleID, 
			memsize, cudaMemcpyHostToDevice) );
  CUDA_CALL( cudaMemcpy(_DevVY, sim.GetVY() + _MinParticleID, 
			memsize, cudaMemcpyHostToDevice) );
  CUDA_CALL( cudaMemcpy(_DevVM, sim.GetVM() + _MinParticleID, 
			memsize, cudaMemcpyHostToDevice) );
  CUDA_CALL( cudaMemcpy(_DevAX, sim.GetAX() + _MinParticleID, 
			memsize, cudaMemcpyHostToDevice) );
  CUDA_CALL( cudaMemcpy(_DevAY, sim.GetAY() + _MinParticleID, 
			memsize, cudaMemcpyHostToDevice) );
}

void DeviceController::SyncPartsDeviceToHost()
{
  this->SyncPosDeviceToHost();

  Simulator& sim = Simulator::Instance();
  size_t memsize = (_MaxParticleID - _MinParticleID + 1) * sizeof(DoF);

  CUDA_CALL( cudaMemcpy(sim.GetVX() + _MinParticleID, 
			_DevVX, memsize, cudaMemcpyDeviceToHost) );
  CUDA_CALL( cudaMemcpy(sim.GetVY() + _MinParticleID, 
			_DevVY, memsize, cudaMemcpyDeviceToHost) );
  CUDA_CALL( cudaMemcpy(sim.GetVM() + _MinParticleID, 
			_DevVM, memsize, cudaMemcpyDeviceToHost) );
  CUDA_CALL( cudaMemcpy(sim.GetAX() + _MinParticleID, 
			_DevAX, memsize, cudaMemcpyDeviceToHost) );
  CUDA_CALL( cudaMemcpy(sim.GetAY() + _MinParticleID, 
			_DevAY, memsize, cudaMemcpyDeviceToHost) );
}

void DeviceController::SyncPosHostToDevice()
{
  Simulator& sim = Simulator::Instance();
  size_t memsize = sim.GetParameters().N * sizeof(DoF);

  CUDA_CALL( cudaMemcpy(_DevX, sim.GetX(), memsize, cudaMemcpyHostToDevice) );
  CUDA_CALL( cudaMemcpy(_DevY, sim.GetY(), memsize, cudaMemcpyHostToDevice) );
}

void DeviceController::SyncPosDeviceToHost()
{
  Simulator& sim = Simulator::Instance();
  size_t memsize = (_MaxParticleID - _MinParticleID + 1) * sizeof(DoF);

  CUDA_CALL( cudaMemcpy(sim.GetX() + _MinParticleID, 
			_DevX + _MinParticleID, 
			memsize, cudaMemcpyDeviceToHost) );
  CUDA_CALL( cudaMemcpy(sim.GetY() + _MinParticleID, 
			_DevY + _MinParticleID, 
			memsize, cudaMemcpyDeviceToHost) );
}
